iT邦幫忙

2025 iThome 鐵人賽

DAY 20
0
Software Development

渲染與GPU編程系列 第 20

Day 19|WebGPU Compute Shader 的基礎範例

  • 分享至 

  • xImage
  •  

上一篇我們對 Compute Shader有了基本的認識 ,今天把觀念補齊:
工作群組(Workgroup) 怎麼分工?同步(Barrier) 在避免資料競爭扮演什麼角色?共享記憶體(var<workgroup> 為何能讓效能暴衝?最後給你一份效能小抄,讓你的 WebGPU 計算從「能跑」變成「跑很快」。


1)三種身分:每個 GPU 執行緒的座標

先用一個好懂的「工地」比喻:

  • Workgroup(工作群組):一隊工人。你用 @workgroup_size(x,y,z) 決定「每隊幾個人」。
  • Workgroup ID@builtin(workgroup_id)):第幾隊。
  • Local Invocation ID@builtin(local_invocation_id)):隊內第幾號工人。
  • Global Invocation ID@builtin(global_invocation_id)):把「第幾隊 × 每隊人數 + 隊內編號」合起來的全場編號。常用來對資料做索引。

口訣:幾隊(dispatch) × 每隊人數(workgroup_size) = 這次總執行緒數

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id)  gid  : vec3<u32>,
        @builtin(workgroup_id)          g    : vec3<u32>,
        @builtin(local_invocation_id)   lid  : vec3<u32>) {
  // gid.x   → 這個人要處理的「全域」資料索引
  // g.x     → 我是第幾隊
  // lid.x   → 我在隊內的編號(0..63)
}

2)同步是什麼?為什麼一定要用 Barrier

多個執行緒「一起寫同一塊資料」會產生資料競爭(race condition)
Barrier(柵欄) 讓所有執行緒在關鍵點「等彼此」,並且保證前面的寫入,後面看得到

在 WGSL 裡有兩種柵欄(只在同一個 workgroup 之內有效):

  • workgroupBarrier()執行緒同步 + workgroup 記憶體var<workgroup>)的可見性同步。

  • storageBarrier():針對 Storage(SBO) 的可見性同步。

    注意:沒有跨 workgroup 的 Barrier。如果你需要「全場一起同步」,做法是分多個 pass(先寫、再讀)。


3)共享記憶體 var<workgroup>:GPU 上的「超快暫存區」

  • 讀寫延遲比 Storage Buffer 小很多(靠近計算核心)。
  • 典型用法:大家先把資料載到 workgroup 記憶體 → workgroupBarrier() 等齊 → 在這塊快取裡做多輪運算 → 只把結果寫回 Storage
  • 限制:容量不大(看 device.limits.maxComputeWorkgroupStorageSize,常見 ~32KB)。陣列大小必須是常數

4) 基礎範例

在CPU端建立一個陣列存放1~10,並塞入GPU buffer中。
執行compute shader,將每個元素乘以2後寫回。

async function initWebGPU() {
    if (!navigator.gpu) {
        console.error("WebGPU is not supported in this browser.");
        return;
    }

    const adapter = await navigator.gpu.requestAdapter();
    const device = await adapter.requestDevice();
    const canvas = document.querySelector('canvas');    
    const context = canvas.getContext('webgpu');

    const shaderCode = `
		@group(0) @binding(0) var<storage, read_write> data: array<u32>;

        @compute @workgroup_size(1)
        fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
            let index = global_id.x;
        	data[index] = data[index] * 2;
        }
    `;

	//==============================
	//		    Buffer
	//==============================
	const storageBuffer = device.createBuffer({
		size: 4 * 10, // a 32 bit uint
		usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC,
	});
    const initialData = new Uint32Array([1, 2, 3, 4, 5, 6, 7, 8, 9, 10]);
    device.queue.writeBuffer(storageBuffer, 0, initialData);

	//===============================
	//			Binding
	//===============================
	const bindGroupLayout = device.createBindGroupLayout({
		entries: [
			{
				binding: 0,
				visibility: GPUShaderStage.COMPUTE,
				buffer: { type: 'storage' },
			},
		],
	});
	
	const bindGroup = device.createBindGroup({
		layout: bindGroupLayout,
		entries: [
			{
				binding: 0,
				resource: {
					buffer: storageBuffer,
				},
			},
		],
	});
	//===============================
	//			Pipeline
	//===============================
	const shaderModule = device.createShaderModule({
		code: shaderCode,
	});
	
	const pipeline = device.createComputePipeline({
		layout: device.createPipelineLayout({
			bindGroupLayouts: [bindGroupLayout],
		}),
		compute: {
			module: shaderModule,
			entryPoint: "main",
		},
	});
	

    const commandEncoder = device.createCommandEncoder();
    const passEncoder = commandEncoder.beginComputePass();
    passEncoder.setPipeline(pipeline);
	passEncoder.setBindGroup(0, bindGroup);
    passEncoder.dispatchWorkgroups(10, 1, 1);
    passEncoder.end();

    const commandBuffer = commandEncoder.finish();
    device.queue.submit([commandBuffer]);

    console.log("Compute shader executed.");

}

initWebGPU();

CPU讀取資料 (範例)

準備一個COPY_DST 的buffer裝寫出的資料,再用console.log print出來。

    const debug_size = 10;
    const _data_size = 4; // size of int 
    const _buffer_size = debug_size * _data_size;		
    const buffer = storageBuffer;       

    const readBuffer = device.createBuffer({			
        size: _buffer_size,
        usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
        mappedAtCreation: false,
        label: "read back buffer"
    });
    const _cmdPass = device.createCommandEncoder();
    
    _cmdPass.copyBufferToBuffer(buffer, 0, readBuffer, 0, _buffer_size);
    device.queue.submit([ _cmdPass.finish()]);

    console.log("=============== Read Back =================");
    readBuffer.mapAsync(GPUMapMode.READ).then(() => {
        const result = new Uint32Array(readBuffer.getMappedRange());

        for(var i = 0 ; i < debug_size ;i++ ){
            console.log(result[i]);
        }
        
        readBuffer.unmap();
    });

成果:
result


5)Workgroup 大小怎麼選?(跨硬體的穩健預設)

  • 查限制:device.limits.maxComputeWorkgroupSizeX/Y/ZmaxComputeInvocationsPerWorkgroup(常見 256 或 1024)、maxWorkgroupsPerDimensionmaxComputeWorkgroupStorageSize

  • 穩健的起點:

    • 一維資料:@workgroup_size(64)128
    • 二維(影像/矩陣):@workgroup_size(8,8)(16,16)(看 workgroup 記憶體大小)。
  • 調整原則:

    • 盡量讓一個 workgroup 的執行緒數是 32/64 的倍數(多數 GPU 的 SIMD/warp 大小)。
    • coalesced access:相鄰的執行緒讀相鄰的位址。
    • workgroup 記憶體用量 = 元素大小 × tile 面積,要小於 maxComputeWorkgroupStorageSize

在 JS/TS 端印出限制(快速檢查):

console.table({
  maxWGSizeX: device.limits.maxComputeWorkgroupSizeX,
  maxWGSizeY: device.limits.maxComputeWorkgroupSizeY,
  maxWGSizeZ: device.limits.maxComputeWorkgroupSizeZ,
  maxInvocations: device.limits.maxComputeInvocationsPerWorkgroup,
  maxWGStorage: device.limits.maxComputeWorkgroupStorageSize
});

6)效能小抄(踩點版)

  1. 記憶體佈局
  • 讓索引 i = global_id.x,相鄰執行緒讀相鄰元素(coalesced)。
  • 盡量用 vec4<f32> 等 16-byte 對齊的結構(特別在 Storage Buffer)。
  • 避免 vec3(以 vec4 取代或 @align(16))。
  1. 工作分配
  • Workgroup size 選 64/128/256 等「硬體友善」數字;不要太小(啟動成本高),也不要大到超過 maxComputeInvocationsPerWorkgroup
  • 「每人多處理幾筆」能降低迭代次數(像上面歸約一次處理兩筆)。
  1. 共享記憶體與同步
  • 把可重用資料搬到 var<workgroup>,用 workgroupBarrier() 把每輪運算切段。
  • 需要跨 workgroup 的整體同步 → 分成多個 pass
  1. Atomic 的用法
  • 先 local 彙整、最後再原子寫回全域,避免大量競爭。
  • 只在必要欄位用 atomic;其他用普通變數可快很多。
  1. 資源與提交流程
  • Pipeline/BindGroup 建立一次、重用。
  • 每幀只更新需要變動的 buffer,用 queue.writeBuffer
  • 把很多小 dispatch 併成較少次大 dispatch(降低 CPU/Driver 開銷)。
  1. 量測
  • performance.now()queue.submit() + await device.queue.onSubmittedWorkDone() 做粗略時間。
  • 做 A/B 比較時,先熱身幾次(JIT/快取暖機)。

7)除錯與常見錯誤

  • 結果隨機/偶爾錯:少了 workgroupBarrier() 或誤用 storageBarrier();或多執行緒同時寫同一格沒有用 atomic。
  • 報對齊/越界:Storage Buffer 結構對齊不對;或 global_id 沒做越界檢查。
  • 跑很慢:workgroup 太小/太大、記憶體訪問跳躍(非 coalesced)、每幀重建 pipeline。
  • 裝置限制造成錯誤:tile 大小超過 maxComputeWorkgroupStorageSize;一次 dispatch 的維度超過 maxWorkgroupsPerDimension


上一篇
Day 18|Compute Shader on Website:WebGPU 的 Storage Buffer(SBO)與 Compute Shader 入門(零基礎友善版)#1
下一篇
Day 20|在 WebGPU 中載入並繪製模型
系列文
渲染與GPU編程22
圖片
  熱門推薦
圖片
{{ item.channelVendor }} | {{ item.webinarstarted }} |
{{ formatDate(item.duration) }}
直播中

尚未有邦友留言

立即登入留言